.func coreSpinLock(.reg .u64 %lock)
{
	.reg .u32 lastValue;
	.reg .pred p;

$whileBegin:
	// while(InterlockedCompareExchange((__int32*)&lock, 1, 0) != 0)
	atom.global.cas.b32 lastValue, [%lock], 1U, 0U;
	setp.ne.u32 p, lastValue, 0U;
@p	bra $whileBegin;
	ret;
}

.func coreSpinUnlock(.reg .u64 %lock)
{
	// lock = 0
	//st.volatile.global.u64 [%lock], 0U;
	st.global.u32 [%lock], 1U;
	ret;
}
